#ifndef CUFALLOC_H
#define CUFALLOC_H
#define FALLOC_H

/*
 *	This is the header file supporting cuFalloc.cu and defining both the host and device-side interfaces. See that file for some more
 *	explanation and sample use code. See also below for details of the host-side interfaces.
 *
 *  Quick sample code:
 *
	#include "cuFalloc.cu"
 	
	__global__ void TestFalloc(fallocHeap* heap)
	{
		fallocInit(heap);

		// create/free heap
		void* obj = fallocGetChunk(heap);
		fallocFreeChunk(heap, obj);

		// create/free alloc
		fallocContext* ctx = fallocCreateCtx(heap);
		char* testString = (char*)falloc(ctx, 10);
		int* testInteger = falloc<int>(ctx);
		fallocDisposeCtx(ctx);
	}

	int main()
	{
		cudaFallocHost fallocHost = cudaFallocInit(1);

		// test
		TestFalloc<<<1, 1>>>(fallocHost.heap);

		// free and exit
		cudaFallocEnd(fallocHost);
		printf("\ndone.\n"); // char c; scanf("%c", &c);
		return 0;
	}
 */

typedef struct _cuFallocHeap fallocHeap;

#ifndef HOSTONLY

///////////////////////////////////////////////////////////////////////////////
// DEVICE SIDE
// External function definitions for device-side code
__device__ void fallocInit(fallocHeap* heap);
__device__ void* fallocGetChunk(fallocHeap* heap);
__device__ void* fallocGetChunks(fallocHeap* heap, size_t length, size_t* allocLength = nullptr);
__device__ void fallocFreeChunk(fallocHeap* heap, void* obj);
__device__ void fallocFreeChunks(fallocHeap* heap, void* obj);
// ALLOC
typedef struct _cuFallocContext fallocContext;
__device__ fallocContext* fallocCreateCtx(fallocHeap* heap);
__device__ void fallocDisposeCtx(fallocContext* ctx);
__device__ void* falloc(fallocContext* ctx, unsigned short bytes, bool alloc = true);
__device__ void* fallocRetract(fallocContext* ctx, unsigned short bytes);
__device__ void fallocMark(fallocContext* ctx, void* &mark, unsigned short &mark2);
__device__ bool fallocAtMark(fallocContext* ctx, void* mark, unsigned short mark2);
//
template <typename T> __device__ T* falloc(fallocContext* ctx) { return (T*)falloc(ctx, sizeof(T), true); }
template <typename T> __device__ void fallocPush(fallocContext* ctx, T t) { *((T*)falloc(ctx, sizeof(T), false)) = t; }
template <typename T> __device__ T fallocPop(fallocContext* ctx) { return *((T*)fallocRetract(ctx, sizeof(T))); }

#endif // HOSTONLY

#ifndef CLIENTONLY

///////////////////////////////////////////////////////////////////////////////
// HOST SIDE
// External function definitions for host-side code

typedef struct {
	fallocHeap* heap;
	int length;
	void* reserved;
} cudaFallocHost;

//
//	cudaFallocInit
//
//	Call this to initialise a falloc heap. If the buffer size needs to be changed, call cudaFallocEnd()
//	before re-calling cudaFallocInit().
//
//	The default size for the buffer is 1 megabyte. For CUDA
//	architecture 1.1 and above, the buffer is filled linearly and
//	is completely used.
//
//	Arguments:
//		length - Length, in bytes, of total space to reserve (in device global memory) for output.
//
//	Returns:
//		cudaSuccess if all is well.
//
extern "C" cudaFallocHost cudaFallocInit(size_t length=1048576, cudaError_t* error=nullptr, void* reserved=nullptr);   // 1-meg

//
//	cudaFallocEnd
//
//	Cleans up all memories allocated by cudaFallocInit() for a heap.
//	Call this at exit, or before calling cudaFallocInit() again.
//
extern "C" void cudaFallocEnd(cudaFallocHost &host);


#ifdef TRACE
///////////////////////////////////////////////////////////////////////////////
// TRACE: HOST SIDE

typedef struct _cudaFallocTrace cudaFallocTrace;

//
//	cudaFallocWTraceInit
//
//	Call this to initialise a falloc heap. If the buffer size needs to be changed, call cudaFallocEnd()
//	before re-calling cudaFallocInit().
//
//	The default size for the buffer is 1 megabyte. For CUDA
//	architecture 1.1 and above, the buffer is filled linearly and
//	is completely used.
//
//	Arguments:
//		length - Length, in bytes, of total space to reserve (in device global memory) for output.
//
//	Returns:
//		cudaSuccess if all is well.
//
extern "C" cudaFallocHost cudaFallocWTraceInit(size_t length=1048576, cudaError_t* error=nullptr);   // 1-meg

//
//	cudaFallocWTraceEnd
//
//	Cleans up all memories allocated by cudaFallocInit() for a heap.
//	Call this at exit, or before calling cudaFallocInit() again.
//
extern "C" void cudaFallocWTraceEnd(cudaFallocHost &host);

//
//  cuFallocSetTraceInfo
//
//	Sets a trace Info.
//
extern "C" void cudaFallocSetTraceInfo(size_t id, bool showDetail);

//
//  cudaFallocTraceInit
//
//	Creates a trace Stream.
//
extern "C" cudaFallocTrace* cudaFallocTraceInit();

//
//	cudaFallocTraceStream
//
//	Streams till empty.
//
extern "C" void* cudaFallocTraceStream(cudaFallocHost &host, cudaFallocTrace* trace, size_t &length);

//
//  cudaFallocTraceEnd
//
//	Frees a trace Stream.
//
extern "C" void cudaFallocTraceEnd(cudaFallocTrace* trace);
#endif

#endif // CLIENTONLY

#endif // CUFALLOC_H